#include <chrono>
#include <stdio.h>
#include <stdint.h>
#include <fstream>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <cuda/pipeline>
#include <cooperative_groups.h>
#include <mma.h>

using namespace nvcuda;

// #include <cuda/barrier>
// namespace cg = cooperative_groups;
// using barrier = cuda::barrier<cuda::thread_scope_block>;

#define BS 32
#define uchar unsigned char
#define DIV_UP(size, align_size) (((size) + (align_size) - 1) / (align_size))
#define ST(X)  uint64_t start##X = clock();
#define ET(X)  uint64_t end##X = clock();

static std::uint64_t now() {
  return std::chrono::duration_cast<std::chrono::milliseconds>
    (std::chrono::high_resolution_clock::now().time_since_epoch()).count();
}

static void dumpbin(void *cuda_mem, size_t sz, const std::string &path) {
  char * buf = new char[sz];
  cudaMemcpy(buf, cuda_mem, sz, cudaMemcpyDeviceToHost);
  {
    std::ofstream ofs(path, std::ios::out | std::ios::trunc | std::ios::binary);
    if (ofs.is_open()){
      ofs.write((const char *)buf, sz);
      // printf("# dump %s\n", path.c_str());
    }else{
      printf("open %s failed!", path.c_str());
    }
  }
}

static void loadbin(void *cuda_mem, size_t sz, const std::string &path) {
  char * buf = new char[sz];
  {
    std::ifstream ifs(path, std::ios::in | std::ios::binary);
    if (ifs.is_open()){
      ifs.read((char *)buf, sz);
      cudaMemcpy(cuda_mem, buf, sz, cudaMemcpyHostToDevice);
      // printf("# load %s\n", path.c_str());
    }else{
      printf("open %s failed!", path.c_str());
    }
  }
}

template<typename T>
__global__ void mem_set(T *src, int w, T value, T bias){
  const int bx = blockIdx.x;
  const int tx = threadIdx.x;
  const int bw = blockDim.x;
  
  int idx = bx*bw + tx;
  if(idx < w){
    if(idx % 2){
      src[idx] = value+bias;
    }else{
      src[idx] = value-bias;
    }
  }
}

struct Case{
  int m;
  int n;
  int k;
  int loop;
  float avg_latency;
};